Cuda Path Tracer开发 [ Bingo ]

Cuda Path Tracer开发

选择smallpt了作为Cuda Path Tracer初尝的借鉴，因此此次主要任务是对smallpt进行Cuda移植，并与好友shawnlu尝试性的做了部分GPU上的优化。代码已开源至smallptCuda

Cuda部分的入门教程较多，可以参考CUDA C/C++ Basics以及An Even Easier Introduction to CUDA

Benchmark

Config	GTX1080Ti	Intel Xeon E5 (6C12T) 2.80GHz
Resolution	1024*768	1024*768
SPP	5000	5000
Cost Time	4.3s	32min

Config	GTX750	Intel Xeon E5 (8C16T) 2.40GHz
Resolution	768*768	768*768
SPP	2048	2048
Cost Time	19.0s	7.2min

Cuda部分杂记

Cuda在Windows上会遇到Kernel执行时间过长而被结束进程的问题，可以参考Note中对核函数的监控项进行修改的操作。
本例使用的helper_math.h可以在Custom Cuda Dir\CUDA Samples\v8.0\common\inc
中找到，Cuda提供了基本的float3, int3之类的包装便于计算。

从OpenMp到Cuda

Cuda的部分难度主要是对于Grid, Block, Thread的理解上, 对于GPU与Cuda的抽象逻辑对应上可以参考Cuda的Threading: Block和Grid的设定与Warp。

总结一下，Warp是组成线程组的基本单位，一般为32，不足32的会以32个线程打包运行(有部分线程不工作，造成浪费)，一般而言BlockSize设定为256总体效果都还不错。SM会根据分配的Block消耗资源的多少来分配究竟有多少个Block在SM中。

具体的N卡核心数等属性信息，可自行获取cudaDeviceProp打印。具体属性值含义可参考cudaDeviceProp Struct Reference

和OpenMp这类CPU上的多核并行不同的是，在线程中并不确定执行顺序，因此需要知晓当前线程对应像素关系就需要threadIdx, blockIdx, blockDim, gridDim等变量来算出。这部分在CUDA C/C++ Basics有简介。

失败的优化

在参考Yining Karl Li的报告后，曾尝试转变原有的基于像素的并行转为基于光线的并行，核心思想就是通过预分配光线，将光线的深度iteration拆分为广度每循环一次加深一层并做stream compaction减少光线数量。

可查看dev分支的历史commit查看具体实现

这里遇到了两个问题

未预估到光线的占用内存情况，按照下方的Ray布局一根光线将会占用60B, 取SPP为2048为例每个像素将预分配SPP根光线，需要 $1024$ 的显存。

struct Ray
{
    float3 origin;
    float3 direction;

    int depth, pixelIndex;
    float3 throughput, L;
};

即便是最简形式的Ray，也仍然占用12B，上述分辨率下也需要占用9G显存，GTX750上的2G显存远远不够(好友的GTX1080ti 12G显存面对45G这样的压力也完全不够)。

struct Ray
{
    float3 origin, direction;
};

2.Cuda中Thrust的Stream Compaction本身并行效率很可观(优化原理可参考Yining Karl Li的报告)，但此法其实更适合于实时光线跟踪，对于目前的GPU纯计算而言，不仅仅增加了Device与Host之间的带宽开销，增加了多次Kernel启动消耗，更增加了线程需要从Global Memory多次读写数据的开销(且存在线程竞争情况)。甚至降低的光线数量带来的性能优化完全不足以弥补。

// cast all rays we need
initRays<<<A, B>>>();

// Russian Roulette would stop ray
while(1)
{
    rayStep<<<C, D>>>();

    // reduce the count of ray
    streamCompaction();
}

最后dev版本放弃了Stream Compaction，改用如下方式，但初始化分配光线仍然需要巨大资源消耗，且在ray结束迭代需要将结果存入放在Global Memory上的output缓冲区时，需要做原子操作以免出现竞争问题(进一步带来了消耗)。

// cast all rays we need
initRays<<<A, B>>>();

// Russian Roulette would stop ray
rayIterate<<<C, D>>>();

由于相比于基于像素并行带来的额外内存开销，整体性能倒退2x左右。原本只是希望展开GPU线程中忙忙的for(spp)繁重的循环已达到性能优化，但基于光线的并行的确更适合实时光线追踪而非纯GPU渲染。

内存上的优化

并行效率难以优化的前提下，内存优化是突破口。由于GPU的Cache过小，不会做类似CPU上的多级Cache块调入优化，Global Memory类似于CPU上的内存，Shared Memory便可充当Register或Cache的地位(Shared的速度最优情况下接近寄存器)，Global与Shared区别可参考Stackoverflow。

这里做的本质上就是手工模拟Cache与内存块交换思想。以如下输出结果至缓冲为例，

__shared__ float3 temp;
temp = // do some assignment;
// output on global memory
output[i] = temp;

先将结果存入shared上的temp，再转入global上的output，思想类似于多级cache(其余sheard部分优化类似)。还对基本的类型做了内存对齐。但由于光线跟踪带宽要求并不高，所以最后结果优化并不明显，不过仍然有1s左右的提升。

GTX750	Before	After
Resolution	768*768	768*768
SPP	2048	2048
Cost Time	20.517145s	19.000868s